#pragma once
//for the error : identifier "IUnknown" is undefined
#ifdef _WIN32
#define WIN32_LEAN_AND_MEAN
#endif
#include "book.h"
// #include "cpu_bitmap.h"
#include <iostream>

#define PI 3.141592654
// const float PI = 3.141592654;
const int DIR_X=1;
const int DIR_Y=0;
const int BLOCKS=10;
const int THREADS=64;

texture<float, cudaTextureType2D, cudaReadModeElementType> texGPU;
texture<float, cudaTextureType2D, cudaReadModeElementType> tex_X_Sobel;
texture<float, cudaTextureType2D, cudaReadModeElementType> tex_Y_Sobel;
/************************************************************************/
/* for dirY:
	-1  -2  -1
	*	*	*
	1	2	1

	for dirX:
	-1	*	1
	-2	*	2
	-1	*	1
	*/

/************************************************************************/
__global__ void GPU_Sobel( float* out_img_buffer, int dir)
{
	int x = threadIdx.x + blockIdx.x*blockDim.x;
	int y = threadIdx.y + blockIdx.y*blockDim.y;

	if ( x>0 && x<(gridDim.x-1) && y>0 && y<(gridDim.y-1) )
	{
		float sobVal;
		switch (dir)
		{
		case DIR_X:
			{
				float lef1 = tex2D(texGPU, x-1, y-1);
				float lef2 = tex2D(texGPU, x-1, y);
				float lef3 = tex2D(texGPU, x-1, y+1);
				float rig1 = tex2D(texGPU, x+1, y-1);
				float rig2 = tex2D(texGPU, x+1, y);
				float rig3 = tex2D(texGPU, x+1, y+1);
				sobVal = (1*rig1 + 2*rig2 + 1*rig3) - (1*lef1 + 2* lef2 + 1*lef3);
			}
			break;
		case DIR_Y:
			{
				float up1 = tex2D(texGPU, x-1, y-1);
				float up2 = tex2D(texGPU, x,   y-1);
				float up3 = tex2D(texGPU, x+1, y-1);
				float down1 = tex2D(texGPU, x-1, y+1);
				float down2 = tex2D(texGPU, x,   y+1);
				float down3 = tex2D(texGPU, x+1, y+1);
				sobVal = (1*down1 + 2*down2 + 1*down3) - (1*up1 + 2*up2 + 1*up3);
			}
			break;
		default:
			break;
		}// end switch
		out_img_buffer[y*gridDim.x + x] = sobVal;
	}//end if
	else
		out_img_buffer[y*gridDim.x + x] = tex2D(texGPU, x, y);
}

__global__ void ComputeMagAndAngle(float* out_MagBuf, float* out_AngleBuf)
{
	int x = threadIdx.x + blockIdx.x*blockDim.x;
	int y = threadIdx.y + blockIdx.y*blockDim.y;

	if ( x>0 && x<(gridDim.x-1) && y>0 && y<(gridDim.y-1) )
	{
		float x_gradient = tex2D(tex_X_Sobel, x, y);
		float y_gradient = tex2D(tex_Y_Sobel, x, y);
		float mag= std::sqrt(x_gradient*x_gradient + y_gradient*y_gradient);
		out_MagBuf[y*gridDim.x + x] = mag;
		float imme = std::atan(y_gradient/x_gradient);
		if (x_gradient<0)
		{
			imme+= PI;
		}
		else
		{
			if(y_gradient<0)
			{
				imme+=2*PI;
			}
		}
		//convert to degree
		imme*=(180.0/PI);
		out_AngleBuf[y*gridDim.x+x] = imme;
	}
	else
	{
			out_MagBuf[y*gridDim.x+x] = 0.0;
			out_AngleBuf[y*gridDim.x+x] = 0.0;
	}
}



void gpuGradientMagAndAngle(float* imgBuffer, int width, int height,
							float* outMagImage, float* outAngleImage)
{
	//record event to get elapsed time
	cudaEvent_t start, stop;
	HANDLE_ERROR( cudaEventCreate( &start ) );
	HANDLE_ERROR( cudaEventCreate( &stop ) );
	HANDLE_ERROR( cudaEventRecord(start, 0));

	//use cudaArray
	cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
	cudaArray* img_buffer;
	cudaMallocArray(&img_buffer, &desc, width, height);
	cudaMemcpyToArray(img_buffer, 0, 0, imgBuffer, sizeof(float)*width*height, cudaMemcpyHostToDevice  );
	//bind with texture
	cudaBindTextureToArray(texGPU, img_buffer, desc);

	float* x_SobelBuffer, *y_SobelBuffer, *mag_Buffer, *angle_Buffer;
	HANDLE_ERROR( cudaMalloc( (void**)&x_SobelBuffer, sizeof(float)*width*height )  );
	HANDLE_ERROR( cudaMalloc( (void**)&y_SobelBuffer, sizeof(float)*width*height )	);
	HANDLE_ERROR( cudaMalloc( (void**)&mag_Buffer, sizeof(float)*width*height) );
	HANDLE_ERROR( cudaMalloc( (void**)&angle_Buffer, sizeof(float)*width*height));

	//call kernel function
// 	dim3   blocks(width/THREAD_PER_BLOCK, height/THREAD_PER_BLOCK);
	dim3   blocks(BLOCKS,BLOCKS);
	dim3   threads(THREADS, THREADS);
	GPU_Sobel<<<blocks, threads>>>(x_SobelBuffer, DIR_X );
	GPU_Sobel<<<blocks, threads>>>(y_SobelBuffer, DIR_Y );

	//bind the Sobel gradient buffer as texture.
	cudaArray* x_sobel, *y_sobel;
	cudaMallocArray(&x_sobel, &desc, width, height);
	cudaMemcpyToArray(x_sobel, 0, 0, x_SobelBuffer, sizeof(float)*width*height, cudaMemcpyDeviceToDevice  );
	cudaBindTextureToArray(tex_X_Sobel, x_sobel, desc);

	cudaMallocArray(&y_sobel, &desc, width, height);
	cudaMemcpyToArray(y_sobel, 0, 0, y_SobelBuffer, sizeof(float)*width*height, cudaMemcpyDeviceToDevice  );
	cudaBindTextureToArray(tex_Y_Sobel, y_sobel, desc);

	ComputeMagAndAngle<<<blocks, threads>>>(mag_Buffer, angle_Buffer);

	HANDLE_ERROR( cudaMemcpy(outMagImage, mag_Buffer, sizeof(float)*width*height, cudaMemcpyDeviceToHost) );
	HANDLE_ERROR( cudaMemcpy(outAngleImage, angle_Buffer, sizeof(float)*width*height, cudaMemcpyDeviceToHost) );

	HANDLE_ERROR( cudaEventRecord(stop, 0));
	HANDLE_ERROR( cudaEventSynchronize( stop));
	float elapseTime;
	HANDLE_ERROR( cudaEventElapsedTime(&elapseTime, start, stop));
	printf("elapse time: %3.3f ms\n", elapseTime);

	HANDLE_ERROR( cudaFreeArray(img_buffer));
	HANDLE_ERROR( cudaFreeArray(x_sobel));
	HANDLE_ERROR( cudaFreeArray(y_sobel));
	HANDLE_ERROR( cudaFree(x_SobelBuffer) );
	HANDLE_ERROR( cudaFree(y_SobelBuffer) );
	HANDLE_ERROR( cudaFree(mag_Buffer));
	HANDLE_ERROR( cudaFree(angle_Buffer));
	cudaUnbindTexture(texGPU);
	cudaUnbindTexture(tex_X_Sobel);
	cudaUnbindTexture(tex_Y_Sobel);
}